
#include <stdio.h>
#define BLOCK_SIZE 32

__global__ void exclusive_sum(int* input, int* output, int N) {
    extern __shared__ int temp[];
    int tid = threadIdx.x;
    int idx = blockIdx.x * blockDim.x + tid;
    
    if (idx < N) {
        temp[tid] = input[idx];
    } else {
        temp[tid] = 0;
    }
    __syncthreads();

    // Hillis-Steele扫描算法
    for (int stride = 1; stride < blockDim.x; stride <<= 1) {
        if (tid >= stride) {
            float val = temp[idx] + temp[idx - stride];
            __syncthreads();
            temp[idx] = val;
        }
        __syncthreads();
    }

    // 转换为exclusive scan
    output[tid] = (tid == 0 ? 0 : temp[tid-1]);
}

void exclusive_scan(int* h_input, int* h_output, int N) {
    int *d_input, *d_output;
    cudaMalloc(&d_input, N * sizeof(int));
    cudaMalloc(&d_output, N * sizeof(int));
    cudaMemcpy(d_input, h_input, N * sizeof(int), cudaMemcpyHostToDevice);
    
    dim3 block(BLOCK_SIZE);
    dim3 grid((N + BLOCK_SIZE - 1) / BLOCK_SIZE);
    exclusive_sum<<<grid, block, BLOCK_SIZE>>>(d_input, d_output, N);
    
    cudaMemcpy(h_output, d_output, N * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(d_input);
    cudaFree(d_output);
}

int main() {
    const int n = 10;
    int h_input[n] = {1,2,3,4,5,6,7,8,9,10};
    int h_output[n];
    exclusive_scan(h_input, h_output, n);
    for (int i = 0; i < n; i++) {
        printf("%3d%s", h_input[i], ((i % 10) != 9) ? " " : "\n");
    }
    printf("exclusive sum:\n");
    for (int i = 0; i < n; i++) {
        printf("%3d%s", h_output[i], ((i % 10) != 9) ? " " : "\n");
    }
    
    return 0;
}